#ifndef _SVM_KERNEL_
#define _SVM_KERNEL_

#include "reduce.h"
#include <math.h>

//compute phi(A'A)
static __device__ __host__ float selfKernel(float *pointerA, int pitchA, float *pointerAEnd) {
    float acc = 0.0f;
    do {
        float value = *pointerA;
        acc += value*value;
        pointerA += pitchA;
    } while (pointerA < pointerAEnd);
    return acc;
}
    
//compute phi(A'B)
static __device__ __host__ float kernel(float *pointerA, int pitchA, float *pointerAEnd, float *pointerB, int pitchB) {
    float acc = 0.0f;
    do {
        acc += (*pointerA) * (*pointerB);
        pointerA += pitchA;
        pointerB += pitchB;
    } while (pointerA < pointerAEnd);
    return acc;
}

//compute phi(A'B) in parallel
static __device__ void parallelKernel(float *pointerA, float *pointerAEnd, float *pointerB, float *sharedTemps) {
    pointerA += threadIdx.x;
    pointerB += threadIdx.x;
    sharedTemps[threadIdx.x] = 0.0f;
    while (pointerA < pointerAEnd) {
        sharedTemps[threadIdx.x] += (*pointerA) * (*pointerB);
        pointerA += blockDim.x;
        pointerB += blockDim.x;
    }
    __syncthreads();

    sumReduce(sharedTemps);
}

//compute phi(A'B) and phi(A'C), where B and C are already in shared memory
//and A is in global memory
static __device__ __host__ void dualKernel(float *pointerA, int pitchA, float *pointerAEnd, float *pointerB, int pitchB, float *pointerC, int pitchC, float& phi1, float& phi2) {
    float acc1 = 0.0f; float acc2 = 0.0f;
    do {
        float a = *pointerA;
        acc1 += a * (*pointerB);
        acc2 += a * (*pointerC);
        pointerA += pitchA;
        pointerB += pitchB;
        pointerC += pitchC;
    } while (pointerA < pointerAEnd);
    phi1 = acc1; phi2 = acc2;
}

#endif
